#include "AccumulateRadiance.cuh"

using namespace OBR;

const uint __accumulate_batch_size = 64;

__global__ void kernel(Spectrum *new_fbo, Spectrum *acc_fbo, size_t accumulated,
                       size_t n_new_samples, size_t n_total)
{
    float ratio = (float)accumulated / (float)n_new_samples;
    float div = 1.0 + ratio;

    uint idx = threadIdx.x + blockIdx.x * blockDim.x;
    for (size_t i = idx * __accumulate_batch_size;
         i < idx * __accumulate_batch_size + __accumulate_batch_size && i < n_total; i++)
    {
        acc_fbo[i] *= ratio;
        acc_fbo[i] += new_fbo[i];
        acc_fbo[i] /= div;
    }
}

void PostProcess::accumulate_radiance_cuda(size_t w, size_t h, size_t d_new_fbo, size_t d_acc_fbo,
                                           size_t accumulated, size_t n_new_samples)
{
    int n_total = w * h;
    int block_size;    // The launch configurator returned block size
    int grid_size;     // The actual grid size needed, based on input size
    int min_grid_size; // The minimum grid size needed to achieve the maximum occupancy

    cudaOccupancyMaxPotentialBlockSize(&min_grid_size, &block_size, kernel);
    // Round up according to array size
    grid_size = (n_total + (block_size * __accumulate_batch_size) - 1) /
                (block_size * __accumulate_batch_size);
    grid_size = std::max(grid_size, min_grid_size);

    // call cuda kernel
    kernel<<<grid_size, block_size>>>((Spectrum *)d_new_fbo, (Spectrum *)d_acc_fbo, accumulated,
                                      n_new_samples, n_total);
}